OpenCL快速入门教程
原文地址：http://opencl.codeplex.com/wikipage?title=OpenCL%20Tutorials%20-%201
翻译日期：2012年6月4日星期一
这是第一篇真正的OpenCL教程。这篇文章不会从GPU结构的技术概念和性能指标入手。我们将会从OpenCL的基础API开始，使用一个小的kernel作为例子来讲解基本的计算管理。
首先我们需要明白的是，OpenCL程序是分成两部分的：一部分是在设备上执行的（对于我们，是GPU），另一部分是在主机上运行的（对于我们，是CPU）。在设备上执行的程序或许是你比较关注的。它是OpenCL产生神奇力量的地方。为了能在设备上执行代码，程序员需要写一个特殊的函数（kernel函数）。这个函数需要使用OpenCL语言编写。OpenCL语言采用了C语言的一部分加上一些约束、关键字和数据类型。在主机上运行的程序提供了API，所以i可以管理你在设备上运行的程序。主机程序可以用C或者C++编写，它控制OpenCL的环境（上下文，指令队列…）。
设备（Device）我们来简单的说一下设备。设备，像上文介绍的一样，OpenCL编程最给力的地方。
我们必须了解一些基本概念：
Kernel：你可以把它想像成一个可以在设备上执行的函数。当然也会有其他可以在设备上执行的函数，但是他们之间是有一些区别的。Kernel是设备程序执行的入口点。换言之，Kernel是唯一可以从主机上调用执行的函数。
现在的问题是：我们如何来编写一个Kernel？在Kernel中如何表达并行性？它的执行模型是怎样的？解决这些问题，我们需要引入下面的概念：
SIMT：单指令多线程（SINGLE INSTRUCTION MULTI THREAD）的简写。就像这名字一样，相同的代码在不同线程中并行执行，每个线程使用不同的数据来执行同一段代码。
Work-item（工作项）：Work-item与CUDA Threads是一样的，是最小的执行单元。每次一个Kernel开始执行，很多（程序员定义数量）的Work-item就开始运行，每个都执行同样的代码。每个work-item有一个ID，这个ID在kernel中是可以访问的，每个运行在work-item上的kernel通过这个ID来找出work-item需要处理的数据。
Work-group（工作组）：work-group的存在是为了允许work-item之间的通信和协作。它反映出work-item的组织形式（work-group是以N维网格形式组织的，N=1，2或3）。
Work-group等价于CUDA thread blocks。像work-items一样，work-groups也有一个kernel可以读取的唯一的ID。
ND-Range:ND-Range是下一个组织级别，定义了work-group的组织形式（ND-Rang以N维网格形式组织的，N=1，2或3）；
这是ND-Range组织形式的例子
Kernel现在该写我们的第一个kernel了。我们写一个小的kernel将两个向量相加。这个kernel需要四个参数：两个要相加的向量，一个存储结果的向量，和向量个数。如果你写一个程序在cpu上解决这个问题，将会是下面这个样子：
void vector_add_cpu (const float* src_a,const float* src_b,float* res,const int num){for (int i = 0; i < num; i++) res[i] = src_a[i] + src_b[i];}
在GPU上，逻辑就会有一些不同。我们使每个线程计算一个元素的方法来代替cpu程序中的循环计算。每个线程的index与要计算的向量的index相同。我们来看一下代码实现：
__kernel void vector_add_gpu (__global const float* src_a, __global const float* src_b, __global float* res,const int num){/* get_global_id(0) 返回正在执行的这个线程的ID。许多线程会在同一时间开始执行同一个kernel，每个线程都会收到一个不同的ID，所以必然会执行一个不同的计算。*/const int idx = get_global_id(0);/* 每个work-item都会检查自己的id是否在向量数组的区间内。如果在，work-item就会执行相应的计算。*/if (idx